
#include <vector_types.h>
#include <math_functions.h>

struct Particula {
    float3 velocidade;
    float massaEspecifica;
    float pressao;
    float3 forcaPressao;
    float3 forcaVisc;
    float3 forcaAceleracao;
};

__global__ void ViscKernel(float3* posicoesParticulasT, int numeroParticulas, int resolucaoGridX, int resolucaoGridY, unsigned int* values, float deltaXGrid, float deltaYGrid, unsigned int* offSet, unsigned int* count, float h, Particula* particulasT, float massa, float coeficienteViscosidade) {
    int thread = threadIdx.x;
    int block = blockIdx.x;
    int dim = blockDim.x;
    int indiceValue = dim * block + thread;
    int jAtual = 0;
    int iAtual = 0;
    int jMin = 0;
    int jMax = resolucaoGridX;
    int iMin = 0;
    int iMax = resolucaoGridY;
    int idCell;
    int kInicial;
    int kFinal;

    if (indiceValue < numeroParticulas) {
        int indiceParticulaAtual = values[indiceValue];
        //forcaViscosidade[indiceParticulaAtual].x = 0;
        //forcaViscosidade[indiceParticulaAtual].y = 0;
        //forcaViscosidade[indiceParticulaAtual].z = 0;
        jAtual = (posicoesParticulasT[indiceParticulaAtual].x / deltaXGrid);
        iAtual = (posicoesParticulasT[indiceParticulaAtual].y / deltaYGrid);

        //VERIFICAO DAS EXTREMIDADES.
        if (jAtual <= 0) {
            jMin = 0;
        } else {
            jMin = jAtual - 1;
        }
        if (jAtual >= (resolucaoGridX - 1)) {
            jMax = jAtual;
        } else {
            jMax = jAtual + 1;
        }
        if (iAtual <= 0) {
            iMin = 0;
        } else {
            iMin = iAtual - 1;
        }
        if (iAtual >= (resolucaoGridY - 1)) {
            iMax = (resolucaoGridY - 1);
        } else {
            iMax = iAtual + 1;
        }

        particulasT[indiceParticulaAtual].forcaVisc.x  = 0;
        particulasT[indiceParticulaAtual].forcaVisc.y  = 0;
        particulasT[indiceParticulaAtual].forcaVisc.z  = 0;
        //CALCULA A FORÇA VISCOSIDADE.
        for (int i = iMin; i <= iMax; i++) {
            for (int j = jMin; j <= jMax; j++) {
                float3 forcaViscTemp;
                forcaViscTemp.x = 0;
                forcaViscTemp.y = 0;
                forcaViscTemp.z = 0;
                idCell = (i * resolucaoGridX) + j;

                kInicial = offSet[idCell];
                kFinal = offSet[idCell] + count[idCell];

                for (int k = kInicial; k < kFinal; k++) {
                    int indiceParticulaVizinha = values[k];
                    if (indiceParticulaVizinha != indiceParticulaAtual) {
                        float distancia = sqrtf((posicoesParticulasT[indiceParticulaAtual].x - posicoesParticulasT[indiceParticulaVizinha].x) * (posicoesParticulasT[indiceParticulaAtual].x - posicoesParticulasT[indiceParticulaVizinha].x) + (posicoesParticulasT[indiceParticulaAtual].y - posicoesParticulasT[indiceParticulaVizinha].y) * (posicoesParticulasT[indiceParticulaAtual].y - posicoesParticulasT[indiceParticulaVizinha].y));
                        if (distancia > 0 && distancia < h) {
                            float kernelVisc = (45 / (3.14 * powf(h, 6))) * (h - distancia);
                            forcaViscTemp.x += 0 * massa * ((particulasT[indiceParticulaVizinha].velocidade.x - particulasT[indiceParticulaAtual].velocidade.x) / 2 * particulasT[indiceParticulaVizinha].massaEspecifica) * kernelVisc;
                            forcaViscTemp.y += 0 * massa * ((particulasT[indiceParticulaVizinha].velocidade.y - particulasT[indiceParticulaAtual].velocidade.y) / 2 * particulasT[indiceParticulaVizinha].massaEspecifica) * kernelVisc;
                        }
                    }
                }
                particulasT[indiceParticulaAtual].forcaVisc.x += forcaViscTemp.x;
                particulasT[indiceParticulaAtual].forcaVisc.y += forcaViscTemp.y;
                particulasT[indiceParticulaAtual].forcaVisc.z += forcaViscTemp.z;
            }
        }
        particulasT[indiceParticulaAtual].forcaVisc.x *= coeficienteViscosidade;
        particulasT[indiceParticulaAtual].forcaVisc.y *= coeficienteViscosidade;
        particulasT[indiceParticulaAtual].forcaVisc.z *= coeficienteViscosidade;
    }
}

extern "C" void launch_ViscKernel(unsigned int blocksPerGrid, unsigned int threadsPerBlock, float3* posicoesParticulasT, int numeroParticulas, int resolucaoGridX, int resolucaoGridY, unsigned int* values, float deltaXGrid, float deltaYGrid, unsigned int* offSet, unsigned int* count, float h, Particula* particulasT, float massa, float coeficienteViscosidade) {
    // execute the kernel
    ViscKernel << <blocksPerGrid, threadsPerBlock >> >(posicoesParticulasT, numeroParticulas, resolucaoGridX, resolucaoGridY, values, deltaXGrid, deltaYGrid, offSet, count, h, particulasT, massa, coeficienteViscosidade);
}

